//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : boundarypsi.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef BOUNDARYPSI_H
#define BOUNDARYPSI_H

#include "kernels/boundarypsi_kernels.h"

namespace wavescomputer
{

namespace boundarypsi
{

namespace reflectiveneumann
{

template <class T>
inline void north(T* psi, 
                  const T* psiN, 
                  const T* psi_in_N_ext, const T* psi_in_N_cen,
                  const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
        
    boundarypsi::reflectiveneumann::kernel_bp_rn_north<T> 
       <<< numBlocks, NUM_THREADS >>> 
          (psi, psiN, psi_in_N_ext, psi_in_N_cen, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::reflectiveneumann::kernel_bp_rn_north");
}

template <class T>
inline void east(T* psi, 
                 const T* psiE, 
                 const T* psi_in_E_ext, const T* psi_in_E_cen,
                 const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx2 / NUM_THREADS);
        
    boundarypsi::reflectiveneumann::kernel_bp_rn_east<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiE, psi_in_E_ext, psi_in_E_cen, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::reflectiveneumann::kernel_bp_rn_east");
}

template <class T>
inline void south(T* psi, 
                  const T* psiS, 
                  const T* psi_in_S_ext, const T* psi_in_S_cen,
                  const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
        
    boundarypsi::reflectiveneumann::kernel_bp_rn_south<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiS, psi_in_S_ext, psi_in_S_cen, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::reflectiveneumann::kernel_bp_rn_south");
}

template <class T>
inline void west(T* psi, 
                 const T* psiW, 
                 const T* psi_in_W_ext, const T* psi_in_W_cen,
                 const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx2 / NUM_THREADS);
        
    boundarypsi::reflectiveneumann::kernel_bp_rn_west<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiW, psi_in_W_ext, psi_in_W_cen, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::reflectiveneumann::kernel_bp_rn_west");
}

} // namespace reflectiveneumann




namespace weaklyreflective
{

template <class T>
inline void north(T* psi, 
                  const T* psiN, 
                  const T* psi_in_N_ext,
                  const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
        
    boundarypsi::weaklyreflective::kernel_bp_wr_north<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiN, psi_in_N_ext, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::weaklyreflective::kernel_bp_wr_north");
}

template <class T>
inline void east(T* psi, 
                 const T* psiE, 
                 const T* psi_in_E_ext,
                 const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx2 / NUM_THREADS);
        
    boundarypsi::weaklyreflective::kernel_bp_wr_east<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiE, psi_in_E_ext, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::weaklyreflective::kernel_bp_wr_east");
}

template <class T>
inline void south(T* psi, 
                  const T* psiS, 
                  const T* psi_in_S_ext,
                  const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx1 / NUM_THREADS);
        
    boundarypsi::weaklyreflective::kernel_bp_wr_south<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiS, psi_in_S_ext, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::weaklyreflective::kernel_bp_wr_south");
}

template <class T>
inline void west(T* psi, 
                 const T* psiW, 
                 const T* psi_in_W_ext,
                 const int ld, const int Nx1, const int Nx2)
{
    int numBlocks = (int)ceil((float)Nx2 / NUM_THREADS);
        
    boundarypsi::weaklyreflective::kernel_bp_wr_west<T> 
        <<< numBlocks, NUM_THREADS >>> 
          (psi, psiW, psi_in_W_ext, ld, Nx1, Nx2);

    checkCudaError(cudaGetLastError(), 
        "wavescomputer::boundarypsi::weaklyreflective::kernel_bp_wr_west");
}

} // namespace weaklyreflective

} // namespace boundarypsi

} // namespace wavescomputer

#endif // BOUNDARYPSI_H

